#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

#define RED 0
#define GREEN 1
#define BLUE 2

__global__ void smooth(unsigned char ***matrix_in, unsigned char ***matrix_out, int width, int height){
	int k, l;
	unsigned char value = 0;
	
	int color = blockIdx.x;
	int j = threadIdx.x;
	int i = threadIdx.y;

	// Applying smooth function
	if((2 <= i) && (i < height-2) && (2 <= j) && (j < width-2)){
		for(k=-2; k<=2; k++){
			for(l=-2; l<=2; l++){
				value += matrix_in[i+k][j+l][color]*0.04;

			}
		}
		matrix_out[i][j][color] = value;
	}
}

int main(int argc, char* argv[]){
	char magic[3];
	char *input = argv[1], *output = argv[2];
	FILE *img;
	int i, j, width, height, max, red, green, blue;
	unsigned char ***mat, ***dev_mat_in, ***dev_mat_out;
	clock_t execution_time;



	// 1. Carregar imagem
	img = fopen(input, "rb");
	fscanf(img, "%s", magic);
	fscanf(img, " %d %d", &width, &height);
	fscanf(img, "%d", &max);
	printf("Type: %s\n", magic);
	printf("%dx%d\n", width, height);
	mat = (unsigned char***)malloc(height*sizeof(unsigned char**));
	for(i=0; i<height; i++){
		mat[i] = (unsigned char**)malloc(width*sizeof(unsigned char*));
		if(!strcmp(magic, "P6")){	// Binary code
			for(j=0; j<width; j++){
				mat[i][j] = (unsigned char*)malloc(3*sizeof(unsigned char));
				fread(mat[i][j], 1, 3, img);
			}
		}
		else if(!strcmp(magic, "P3")){	// ASCII code
			for(j=0; j<width; j++){
				mat[i][j] = (unsigned char*)malloc(3*sizeof(unsigned char));
				fscanf(img, "%d %d %d", &red, &green, &blue);
				mat[i][j][RED] = (unsigned char)(red);
				mat[i][j][GREEN] = (unsigned char)(green);
				mat[i][j][BLUE] = (unsigned char)(blue);
			}
		}
		else{	// Erro
			printf("Erro\n");
			for(i=0; i<height; i++){
				for(j=0; j<width; j++){
					free(mat[i][j]);
				}
				free(mat[i]);
				return -1;
			}
			free(mat);
	 	}
	}
	printf("1 ... OK!\n");
	fclose(img);



	// 2. Copiar imagem para GPU
	printf("2.1\n");
	cudaMalloc((void****)&dev_mat_in, height*sizeof(unsigned char**));
	printf("2.2\n");
	cudaMalloc((void****)&dev_mat_out, height*sizeof(unsigned char**));
	for(i=0; i<height; i++){
		printf("2.3\n");
		cudaMalloc(&dev_mat_in[i], height*sizeof(unsigned char*));
		cudaMalloc(&dev_mat_out[i], height*sizeof(unsigned char*));
		for(j=0; j<height; j++){
			cudaMalloc(&dev_mat_in[i][j], 3*sizeof(unsigned char));
			cudaMalloc(&dev_mat_in[i][j], 3*sizeof(unsigned char));
			cudaMemcpy(dev_mat_in[i][j], mat[i][j], 3*sizeof(unsigned char), cudaMemcpyHostToDevice);
			cudaMemcpy(dev_mat_out[i][j], mat[i][j], 3*sizeof(unsigned char), cudaMemcpyHostToDevice);
		}
	}
	printf("2. ... OK!\n");


	// 3. Iniciar a contagem de tempo
	execution_time = clock();



	// 4. Processar o smooth
	dim3 threads_per_block(width, height);
	smooth<<<3, threads_per_block>>>(dev_mat_in, dev_mat_out, width, height);



	// 5. Finalizar a contagem e exibir tempo de processamento
	execution_time = clock() - execution_time;
	printf("%.6f", ((float)execution_time/CLOCKS_PER_SEC));



	// 6. Trazer imagem processada da CPU
	for(i=0; i<height; i++){
		for(j=0; j<width; j++){
			cudaMemcpy(mat[i][j], dev_mat_out[i][j], 3*sizeof(unsigned char), cudaMemcpyDeviceToHost);
		}
	}

	// 7. Salvar nova imagem
	img = fopen(output, "wb");
	fprintf(img, "%s\n", magic);
	fprintf(img, "%d %d\n", height, width);
	fprintf(img, "%d\n", max);
	if(!strcmp(magic, "P6")){	// Binary code
		for(i=0; i<height; i++){
			for(j=0; j<width; j++){
				fwrite(mat[i][j], 1, 3, img);
			}
		}
	}
	else if(!strcmp(magic, "P3")){	// ASCII code
		for(i=0; i<height; i++){
			for(j=0; j<width; j++){
				fprintf(img, "%d %d %d\t", mat[i][j][RED], mat[i][j][GREEN], mat[i][j][BLUE]);
			}
		}
	}
	fclose(img);

	// 8. Liberar memória
	for(i=0; i<height; i++){
		for(j=0; j<width; j++){
			free(mat[i][j]);
			cudaFree(dev_mat_in[i][j]);
			cudaFree(dev_mat_out[i][j]);
		}
		free(mat[i]);
		cudaFree(dev_mat_in[i]);
		cudaFree(dev_mat_out[j]);
	}
	free(mat);
	cudaFree(dev_mat_in);
	cudaFree(dev_mat_out);


	// 9. FIM
	return 0;
}
